#define K_MAX_SHAPE_DIM 0
#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;      //昇腾双buffer技术
#include <math.h>

// 常数定义
constexpr float THRESHOLD_A = 2.5625;
constexpr float THRESHOLD_B = 36974.0;
constexpr float THRESHOLD_PI = 3.1415926535897932384626433832795;  //M_PI

// 泰勒系数计数
constexpr int COS_COUNT = 5;   //5阶三角函数
constexpr int SIN_COUNT = 5; 
constexpr int SN_COUNT = 5;
constexpr int SD_COUNT = 6;
constexpr int CN_COUNT = 5;
constexpr int CD_COUNT = 6;
constexpr int FN_COUNT = 9;
constexpr int FD_COUNT = 10;
constexpr int GN_COUNT = 10;
constexpr int GD_COUNT = 11;

// 系数数组定义
constexpr float COEF_COS[] = {
    // +2.087675e-13,        // x¹⁸ 项: +1/18! ≈ +2.087675e-13
    // -1.147074e-16,        // x¹⁶ 项: -1/16! ≈ -1.147074e-16
    -0.0000000000114707,  // x^14 系数: -1/14! ≈ -1.14707e-11
     0.000000002087675,    // x^12 系数: +1/12! ≈ +2.08768e-9
    -0.000000275573192,    // x^10 系数: -1/10! ≈ -2.75573e-7
     0.000024801587302,    // x^8  系数: +1/8!  ≈ +2.48016e-5
    -0.001388888888889,    // x^6  系数: -1/6!  ≈ -1.38889e-3
     0.041666666666667,    // x^4  系数: +1/4!  ≈ +4.16667e-2
    -0.500000000000000,    // x^2  系数: -1/2!  ≈ -5.00000e-1
    1.000000000000000000   // 常数项: 1
};

constexpr float COEF_SIN[] = {
    // +8.220635e-18,        // x¹⁹ 项: +1/19! ≈ +8.220635e-18
    // -2.811457e-15,        // x¹⁷ 项: -1/17! ≈ -2.811457e-15
    -7.647163e-13,       // x¹⁵ 项: -1/15! ≈ -1/1,307,674,368,000
    +1.605904e-10,        // x¹³ 项: +1/13! ≈ +1/6,227,020,800
    -2.50521083854e-8,    // x¹¹ 项: -1/11! ≈ -1/39,916,800
    +2.75573192239e-6,    // x⁹  项: +1/9!  ≈ +1/362,880
    -0.000198412698412,   // x⁷  项: -1/7!  ≈ -1/5,040
    +0.00833333333333,    // x⁵  项: +1/5!  ≈ +1/120
    -0.1666666666667,     // x³  项: -1/3!  ≈ -1/6
    +1.0000000000000       // x¹  项: +1/1! = +1
};

constexpr float COEF_CN[] = {
    -4.98843114573573548651E-8,
    9.50428062829859605134E-6,
    -6.45191435683965050962E-4,
    1.88843319396703850064E-2,
    -2.05525900955013891793E-1,
    9.99999999999999998822E-1
};

constexpr float COEF_CD[] = {
    3.99982968972495980367E-12,
    9.15439215774657478799E-10,
    1.25001862479598821474E-7,
    1.22262789024179030997E-5,
    8.68029542941784300606E-4,
    4.12142090722199792936E-2,
    1.00000000000000000118E0
};

constexpr float COEF_SN[] = {
    -2.99181919401019853726E3,
    7.08840045257738576863E5,
    -6.29741486205862506537E7,
    2.54890880573376359104E9,
    -4.42979518059697779103E10,
    3.18016297876567817986E11
};

constexpr float COEF_SD[] = {
    1.0,
    2.81376268889994315696E2,
    4.55847810806532581675E4,
    5.17343888770096400730E6,
    4.19320245898111231129E8,
    2.24411795645340920940E10,
    6.07366389490084639049E11
};

constexpr float COEF_FN[] = {
    4.21543555043677546506E-1,
    1.43407919780758885261E-1,
    1.15220955073585758835E-2,
    3.45017939782574027900E-4,
    4.63613749287867322088E-6,
    3.05568983790257605827E-8,
    1.02304514164907233465E-10,
    1.72010743268161828879E-13,
    1.34283276233062758925E-16,
    3.76329711269987889006E-20
};

constexpr float COEF_FD[] = {
    1.0,
    7.51586398353378947175E-1,
    1.16888925859191382142E-1,
    6.44051526508858611005E-3,
    1.55934409164153020873E-4,
    1.84627567348930545870E-6,
    1.12699224763999035261E-8,
    3.60140029589371370404E-11,
    5.88754533621578410010E-14,
    4.52001434074129701496E-17,
    1.25443237090011264384E-20
};

constexpr float COEF_GN[] = {
    5.04442073643383265887E-1,
    1.97102833525523411709E-1,
    1.87648584092575249293E-2,
    6.84079380915393090172E-4,
    1.15138826111884280931E-5,
    9.82852443688422223854E-8,
    4.45344415861750144738E-10,
    1.08268041139020870318E-12,
    1.37555460633261799868E-15,
    8.36354435630677421531E-19,
    1.86958710162783235106E-22
};

constexpr float COEF_GD[] = {
    1.0,
    1.47495759925128324529E0,
    3.37748989120019970451E-1,
    2.53603741420338795122E-2,
    8.14679107184306179049E-4,
    1.27545075667729118702E-5,
    1.04314589657571990585E-7,
    4.60680728146520428211E-10,
    1.10273215066240270757E-12,
    1.38796531259578871258E-15,
    8.39158816283118707363E-19,
    1.86958710162783236342E-22
};


template<typename TYPE_X,typename TYPE_Y>  class KernelFresnelCos {
using T = DTYPE_Y;
public:
    __aicore__ inline KernelFresnelCos() {}

    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,int32_t CoreDataNum, int32_t finalTileNum, int32_t tileDataNum, int32_t TailDataNum, TPipe* pipeIn) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");

        this->tileDataNum = tileDataNum;
        this->coreDataNum = CoreDataNum;
        this->tileNum = finalTileNum;
        this->tailDataNum = TailDataNum;

        xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x, this->coreDataNum);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y, this->coreDataNum);

        pipe = pipeIn;
        pipe->InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(DTYPE_X));
        pipe->InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(DTYPE_Y));
        pipe->InitBuffer(tmp1, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(tmp2, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(tmp3, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(tmp4, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(B_bits, this->tileDataNum * sizeof(uint8_t));   
        pipe->InitBuffer(B_result, this->tileDataNum * sizeof(float));  
        pipe->InitBuffer(B_zero, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(tmp_x4, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_cn_polevl, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(tmp_x2, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_square, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_fn_polevl, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_sj_polevl, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_f, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_g, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_z, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_c, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_s, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_y, this->tileDataNum * sizeof(float));
        pipe->InitBuffer(data_int, this->tileDataNum * sizeof(int));
    }

    __aicore__ inline void Process()
    {
        int32_t loopCount = this->tileNum;   //手动考虑了双buff
        this->processDataNum = this->tileDataNum;
        for (int32_t i = 0; i < loopCount; i++) {
            if (i == this->tileNum - 1) {
                this->processDataNum = this->tailDataNum;
            }
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        } 
    }
private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
        inQueueX.EnQue(xLocal);
    }

    __aicore__ inline void Compute(int32_t progress)
    {
        LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();

        auto p1 = tmp1.Get<float>();  //x
        auto p2 = tmp2.Get<float>();  //-x
        auto p3 = tmp3.Get<float>();  
        auto bits = B_bits.Get<uint8_t>();
        auto result = B_result.Get<float>();
        auto zero = B_zero.Get<float>();
        auto sign = tmp4.Get<float>();  //符号缓存

        /*if (val < 0) {
            sign = -1.0; val = -val;
        }*/
        Duplicate(zero, (float)0.0, this->processDataNum); 
        if constexpr (std::is_same_v<DTYPE_Y, half> || std::is_same_v<DTYPE_Y, bfloat16_t>){
            Cast(p1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
        }
        else if constexpr (std::is_same_v<DTYPE_Y, float>){
            Adds(p1,xLocal,(float)0.0,this->processDataNum); 
        }
        else{}
        Compare(bits, p1, zero, CMPMODE::LT, this->processDataNum);
        Abs(p1, p1, this->processDataNum);
        Duplicate(zero, float(1.0), this->processDataNum); 
        Duplicate(p3, float(-1.0), this->processDataNum); 
        Select(sign, bits, p3, zero, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum); //sign = -1.0

        /*_generic_fresnel_cos_interval*/
        auto p_tmp_x2 = tmp_x2.Get<float>(); 
        auto p_tmp_x4 = tmp_x4.Get<float>();
        auto p_data_cn_polevl = data_cn_polevl.Get<float>(); 
        Mul(p_tmp_x2, p1, p1, this->processDataNum);
        Mul(p_tmp_x4, p_tmp_x2, p_tmp_x2, this->processDataNum);
        Muls(p_data_cn_polevl, p_tmp_x4, COEF_CN[0], this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CN[1],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CN[2],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CN[3],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CN[4],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CN[5],this->processDataNum);
        Mul(result, p1, p_data_cn_polevl, this->processDataNum);
        Muls(p_data_cn_polevl, p_tmp_x4, COEF_CD[0], this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CD[1],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CD[2],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CD[3],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CD[4],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CD[5],this->processDataNum);
        Mul(p_data_cn_polevl, p_data_cn_polevl, p_tmp_x4, this->processDataNum);
        Adds(p_data_cn_polevl, p_data_cn_polevl, COEF_CD[6],this->processDataNum);
        Div(result, result, p_data_cn_polevl, this->processDataNum);

        //p3 = generic_fresnel_asymp(x)
        auto p_data_square = data_square.Get<float>();
        auto p_data_fn_polevl = data_fn_polevl.Get<float>(); 
        auto p_data_sj_polevl = data_sj_polevl.Get<float>(); 
        auto p_data_f = data_f.Get<float>(); 
        auto p_data_g = data_g.Get<float>(); 
        auto p_data_z = data_z.Get<float>(); 
        auto p_data_c = data_c.Get<float>(); 
        auto p_data_s = data_s.Get<float>(); 
        auto p_data_y = data_y.Get<float>(); 
        auto p_data_int = data_int.Get<int>(); 
        
        Duplicate(zero, float(1.0), this->processDataNum); 
        Muls(p2, p_tmp_x2, float(THRESHOLD_PI), this->processDataNum);
        Div(p3,zero,p2,this->processDataNum);
        Mul(p_data_square,p3,p3,this->processDataNum);
        Duplicate(zero, float(THRESHOLD_PI/2), this->processDataNum); 
        Mul(p_data_z,p_tmp_x2,zero,this->processDataNum);

        // double data_fn_polevl = polevl(data_square, COEF_FN, sizeof(COEF_FN)/sizeof(double)-1);
        Muls(p_data_fn_polevl, p_data_square, COEF_FN[0], this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[1],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[2],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[3],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[4],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[5],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[6],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[7],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[8],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FN[9],this->processDataNum);
        Mul(p_data_f,p_data_fn_polevl,p_data_square,this->processDataNum);
        // double data_fd_p1evl = p1evl(data_square, COEF_FD, sizeof(COEF_FD)/sizeof(double)-1);
        Adds(p_data_fn_polevl, p_data_square, COEF_FD[0], this->processDataNum);
        Mul(p_data_fn_polevl, p_data_square, p_data_fn_polevl, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[1],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[2],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[3],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[4],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[5],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[6],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[7],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[8],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[9],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_FD[10],this->processDataNum);
        Div(p_data_f,p_data_f,p_data_fn_polevl,this->processDataNum);
        Duplicate(zero, float(1.0), this->processDataNum); 
        Sub(p_data_f,zero,p_data_f,this->processDataNum);
        // double data_gn_polevl = polevl(data_square, COEF_GN, sizeof(COEF_GN)/sizeof(double)-1);
        Muls(p_data_fn_polevl, p_data_square, COEF_GN[0], this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[1],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[2],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[3],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[4],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[5],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[6],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[7],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[8],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[9],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GN[10],this->processDataNum);
        Mul(p_data_g,p_data_fn_polevl,p3,this->processDataNum);
        // double data_gd_p1evl = p1evl(data_square, COEF_GD, sizeof(COEF_GD)/sizeof(double)-1);
        Adds(p_data_fn_polevl, p_data_square, COEF_GD[0], this->processDataNum);
        Mul(p_data_fn_polevl, p_data_square, p_data_fn_polevl, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[1],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[2],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[3],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[4],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[5],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[6],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[7],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[8],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[9],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[10],this->processDataNum);
        Mul(p_data_fn_polevl, p_data_fn_polevl, p_data_square, this->processDataNum);
        Adds(p_data_fn_polevl, p_data_fn_polevl, COEF_GD[11],this->processDataNum);
        Div(p_data_g,p_data_g,p_data_fn_polevl,this->processDataNum);

        // data_z 处理
        // # Step 1: 标准化到 [0, 2π)
        //     phi = x % (2 * math.pi)
        Duplicate(zero, 2*(float)THRESHOLD_PI, this->processDataNum); 
        Div(zero, p_data_z, zero, this->processDataNum); 
        Cast(p_data_int, zero, RoundMode::CAST_TRUNC, this->processDataNum);
        Cast(zero, p_data_int, RoundMode::CAST_NONE, this->processDataNum);
        Duplicate(p2, 2*(float)THRESHOLD_PI, this->processDataNum); 
        Mul(zero,zero, p2, this->processDataNum); 
        Sub(p_data_z,p_data_z,zero,this->processDataNum); 
        // # Step 2: 映射到 [0, π]
        //     if phi <= math.pi:
        //         return phi
        //     else:
        //         return 2 * math.pi - phi
        Muls(p2,p_data_z,float(-1.0),this->processDataNum); 
        Adds(p2,p2,float(2*THRESHOLD_PI),this->processDataNum); 
        Duplicate(zero, float(THRESHOLD_PI), this->processDataNum); 
        Compare(bits, p_data_z, zero, CMPMODE::LE, this->processDataNum);
        Select(p3, bits, p_data_z, p2, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);
        // double data_c = cos(data_z);
        Mul(p2, p3, p3,this->processDataNum);
        Muls(p_data_sj_polevl, p3, COEF_COS[0], this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[1],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[2],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[3],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[4],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[5],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[6],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        // Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[7],this->processDataNum);
        // Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        // Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_COS[8],this->processDataNum);
        // Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_c, p_data_sj_polevl, COEF_COS[7],this->processDataNum);

        //# Step 3: 映射到 [-π/2, π/2]
        //     if phi <= math.pi/2:
        //         return phi
        //     else if phi <= math.pi/2*3:
        //         return math.pi - phi
        //     else 
        //         return phi - 2 * math.pi
        Muls(p2,p_data_z,float(-1.0),this->processDataNum); 
        Adds(p2,p2,float(THRESHOLD_PI),this->processDataNum); 
        Duplicate(zero, (float)THRESHOLD_PI/(float)2.0, this->processDataNum); 
        Compare(bits, p_data_z, zero, CMPMODE::LE, this->processDataNum);
        Select(p3, bits, p_data_z, p2, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);
        Duplicate(zero, (float)THRESHOLD_PI*(float)1.5, this->processDataNum);
        Adds(p2,p_data_z,(float)THRESHOLD_PI*(float)(-2.0),this->processDataNum); 
        Compare(bits, p_data_z, zero, CMPMODE::LE, this->processDataNum);
        Select(zero, bits, p3, p2, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);
        // double data_s = sin(data_z);
        Mul(p2, zero,zero,this->processDataNum);
        Muls(p_data_sj_polevl, zero, COEF_SIN[0], this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[1],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[2],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[3],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[4],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[5],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[6],this->processDataNum);
        Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        // Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[7],this->processDataNum);
        // Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        // Adds(p_data_sj_polevl, p_data_sj_polevl, COEF_SIN[8],this->processDataNum);
        // Mul(p_data_sj_polevl, p_data_sj_polevl, p2, this->processDataNum);
        Adds(p_data_s, p_data_sj_polevl, COEF_SIN[7],this->processDataNum);
        Mul(p_data_s, p_data_s, zero, this->processDataNum);
        // double data_y = 1.0 / (x * THRESHOLD_PI);
        Duplicate(zero, float(1.0), this->processDataNum);  
        Muls(p3,p1,THRESHOLD_PI,this->processDataNum);
        Div(p_data_y, zero, p3, this->processDataNum);

        // double res_sub = data_f * data_s - data_g * data_c;
        // return 0.5 + data_y * res_sub;
        Mul(p3,p_data_f,p_data_s,this->processDataNum);
        Mul(zero,p_data_g,p_data_c,this->processDataNum);
        Sub(p2,p3,zero,this->processDataNum);
        Mul(p2,p_data_y,p2,this->processDataNum);
        Adds(p3,p2,float(0.5),this->processDataNum);

        //compare with threshold_a
        Duplicate(zero, float(THRESHOLD_A), this->processDataNum); 
        Compare(bits, p_tmp_x2, zero, CMPMODE::LT, this->processDataNum);
        Select(p2, bits, result, p3, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);
        //compare with threshold_b
        Duplicate(zero, float(THRESHOLD_B), this->processDataNum); 
        Compare(bits, p1, zero, CMPMODE::LT, this->processDataNum);
        Duplicate(zero, float(0.5), this->processDataNum); 
        Select(p3, bits, p2, zero, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);      
        // mul sign
        Mul(result, p3, sign, this->processDataNum);
        if constexpr (std::is_same_v<DTYPE_Y, half> || std::is_same_v<DTYPE_Y, bfloat16_t>)
            Cast(yLocal, result, RoundMode::CAST_RINT, this->processDataNum);
        else 
            Adds(yLocal, result,(float)0.0,this->processDataNum);

        outQueueY.EnQue<DTYPE_Y>(yLocal);
        inQueueX.FreeTensor(xLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_Y>();
        DataCopy(yGm[progress * this->tileDataNum], yLocal, this->processDataNum);
        outQueueY.FreeTensor(yLocal);
    }
private:
    TPipe* pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    TBuf<QuePosition::VECCALC> B_result, B_zero, B_bits;
    TBuf<QuePosition::VECCALC> tmp1, tmp2, tmp3, tmp4;
    TBuf<QuePosition::VECCALC> tmp_x4, data_cn_polevl;
    TBuf<QuePosition::VECCALC> tmp_x2, data_square, data_fn_polevl, data_sj_polevl;
    TBuf<QuePosition::VECCALC> data_f, data_g, data_z, data_c, data_s, data_y;
    TBuf<QuePosition::VECCALC> data_int;
    
    GlobalTensor<DTYPE_X> xGm;
    GlobalTensor<DTYPE_Y> yGm;
    int32_t coreDataNum;
    int32_t tileNum;
    int32_t tileDataNum;
    int32_t tailDataNum;
    int32_t processDataNum;
};



extern "C" __global__ __aicore__ void fresnel_cos(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    TPipe pipe;
    KernelFresnelCos<DTYPE_X,DTYPE_Y> op;
    //补充init和process函数调用内容
    op.Init(x, y, tiling_data.CoreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.TailDataNum,&pipe);
    op.Process();
}